#include "OpenCLSetup.hpp"

#define position_t cl_float4
#define velocity_t cl_float4
#define acceleration_t cl_float4
#define mass_t cl_float
#define time_step_t cl_float
#define myInt cl_uint2


#include "FreeImage.h"

#define GAUSSIAN_TIME 0
#define SOBEL_TIME 1
#define HISTOGRAM_TIME 2
#define FIND_COLORS_TIME 3
#define REPLACE_PIXEL_TIME 4

#define NUM_IMAGES 4

bool SaveImage(char *fileName, char *buffer, int width, int height)
{
    FREE_IMAGE_FORMAT format = FreeImage_GetFIFFromFilename(fileName);
    FIBITMAP *image = FreeImage_ConvertFromRawBits((BYTE*)buffer, width,
                        height, width * 4, 32,
                        0xFF000000, 0x00FF0000, 0x0000FF00);
    return (FreeImage_Save(format, image, fileName) == TRUE) ? true : false;
}



size_t RoundUp(int groupSize, int globalSize)
{
    int r = globalSize % groupSize;
    if(r == 0)
    {
     	return globalSize;
    }
    else
    {
     	return globalSize + groupSize - r;
    }
}
//
// main for imageProc. runs a gaussian, then a series of kernels
// to create a colorized image. runs all ops on a series of images
// as listed in a file read in from the command line.
//
int main(int argc, char** argv) {
    OpenCLWrapper w;
    w.enableProfiling = true;

    int event_number[5] ={0,0,0,0,0};

    int errNum = 0;   

   
    try {

	int WORK_GROUP_SIZE = 256;
	int COLORS_PER_CHANNEL = 4;
	if (argc > 1){
	    WORK_GROUP_SIZE = atoi(argv[1]);
	}
	if (argc > 2){
	    COLORS_PER_CHANNEL = atoi(argv[2]);
	}
        w.createContext();
        w.createCommandQueue();
        w.createProgram("parallelImageOps.cl");
        w.createKernel("noop");
	w.createKernel("build_color_histograms");
	w.createKernel("gaussian_filter");
	w.createKernel("get_colors_from_histograms");
	w.createKernel("build_image_from_new_colors");
	w.createKernel("sobel");

	//Uncomment below to see a bug in NVidia's implementation for the card.
	//According to this link:
	//http://www.khronos.org/registry/cl/sdk/1.0/docs/man/xhtml/clGetDeviceInfo.html
	//Since the card return true for CL_DEVICE_IMAGE_SUPPORT it should have a max width
	//of at least 8192.  Our card only has a max width of 4096. 
	//Also, according to CL_DEVICE_IMAGE2D_MAX_HEIGHT the card supports 32768, but loading an image with height 2000 doesn't work
	//we will use a 4000*1000 image  
	/*
	    size_t maxW = 5;
    clGetDeviceInfo(w.device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t),
                    &maxW, NULL);
	cout << "Max Width: " << maxW << endl;
	
	    cl_bool imaging = false;
    clGetDeviceInfo(w.device, CL_DEVICE_IMAGE_SUPPORT, sizeof(cl_bool),
                    &imaging, NULL);
	cout << "Imaging Support: " << imaging << endl;
	*/
	

	int width;
	int height;
	
	// opens the file and reads in each image name into a vector
	std::ifstream myfile;

    	//myfile.open(argv[1]);	
 	myfile.open("images.txt");
	char name[80] = "images/tiled.tif";
	vector<string> allNames;
	string line;                	



    	


	

	

	

		cout << "Filename: " << name << endl;
		cl_mem test_image = LoadImage(w.context, name, width, height);
cout << "The width is: " << width << endl;
cout << "The height is: " << height << endl;
	cl_image_format clImageFormat;
    	clImageFormat.image_channel_order = CL_RGBA;        
    	clImageFormat.image_channel_data_type = CL_UNSIGNED_INT8;
	
	// creates an image for the output of the gaussian filter
	cl_mem output_image = clCreateImage2D(w.context,
                                       CL_MEM_READ_WRITE,
                                       &clImageFormat,
                                       width,
                                       height,
                                       0,
                                       NULL,
                                       &errNum);
	
	cl_sampler my_sampler= clCreateSampler(w.context,
				     CL_FALSE,
				     CL_ADDRESS_CLAMP_TO_EDGE,
			             CL_FILTER_NEAREST,
				     &errNum
				    );
	
	/*****************
  	 GAUSSIAN Kernel 
	 WE CAN DO THIS AS ONE IMAGE, NO ALTERATIONS TO HOT/KERNEL CODE
	******************/
	
	 w.check(clSetKernelArg(w.kernels["gaussian_filter"], 0, sizeof(cl_mem), &test_image ),
                "Error setting kernel argument 0 in gaussian filter.");
	 w.check(clSetKernelArg(w.kernels["gaussian_filter"], 1, sizeof(cl_mem), &output_image ),
                "Error setting kernel argument 1 in gaussian filter.");
	 w.check(clSetKernelArg(w.kernels["gaussian_filter"], 2, sizeof(cl_sampler), &my_sampler ),
                "Error setting kernel argument 2 in gaussian filter");
	 w.check(clSetKernelArg(w.kernels["gaussian_filter"], 3, sizeof(cl_int), &width ),
                "Error setting kernel argument 3 in gaussian filter");
	 w.check(clSetKernelArg(w.kernels["gaussian_filter"], 4, sizeof(cl_int), &height ),
                "Error setting kernel argument 4 in gaussian filter");

	size_t localWorkSize[2] = { 16, 16 };
    	size_t globalWorkSize[2] =  { RoundUp(localWorkSize[0], width),
                                      RoundUp(localWorkSize[1], height) };


	w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["gaussian_filter"], 2, NULL,
                globalWorkSize, localWorkSize, 0, NULL, &w.events[GAUSSIAN_TIME][event_number[GAUSSIAN_TIME]++]),
                "Error queuing kernel for execution.");




	/*****************
  	 SOBEL Kernel 
	WE CAN DO THIS AS ONE IMAGE, NO ALTERATIONS TO HOST/KERNEL CODE
	******************/
	bool *edge_array = (bool*)malloc(sizeof(bool) * width*height);
	for (int i=0;i<width*height;i++){
		edge_array[i] = false;
	} 
	int my_edges = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                sizeof(bool)* width*height, edge_array, &w.errNum));

	 w.check(clSetKernelArg(w.kernels["sobel"], 0, sizeof(cl_mem), &output_image ),
                "Error setting kernel argument 0 in sobel.");
	w.check(clSetKernelArg(w.kernels["sobel"], 1, sizeof(cl_mem), &w.memObjects[my_edges] ),
                "Error setting kernel argument 1 in sobel.");
	 w.check(clSetKernelArg(w.kernels["sobel"], 2, sizeof(cl_mem), &my_sampler ),
                "Error setting kernel argument 2 in sobel.");
	 w.check(clSetKernelArg(w.kernels["sobel"], 3, sizeof(cl_int), &width ),
                "Error setting kernel argument 3 in sobel.");
	 w.check(clSetKernelArg(w.kernels["sobel"], 4, sizeof(cl_int), &height ),
                "Error setting kernel argument 4 in sobel.");


	w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["sobel"], 2, NULL,
                globalWorkSize, localWorkSize, 0, NULL, &w.events[SOBEL_TIME][event_number[SOBEL_TIME]++]),
                "Error queuing kernel for execution.");


	/***********************
  	BUILD HISTOGRAMS Kernel
	WE CANNOT DO THIS AS WE DID IN THE HOST.  ALTERATIONS TO HOST+KERNEL CODE
	************************/
	// create and initialize the histogram matrices to 0	
	int r[256*NUM_IMAGES];
	int g[256*NUM_IMAGES];
	int b[256*NUM_IMAGES];
	int i = 0;
	for(;i<256*NUM_IMAGES;i++){
		r[i] = 0;
		g[i] = 0;
		b[i] = 0;
	}

        
        size_t hist_localWorkSize[1] = { WORK_GROUP_SIZE } ;
	size_t hist_globalWorkSize[1] =  { RoundUp(hist_localWorkSize[0], width*height) } ;


	int single_image_width = width / NUM_IMAGES;
    	int mem_r = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                sizeof(int)* 256*NUM_IMAGES, r, &w.errNum));
 	int mem_g = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                sizeof(int) * 256*NUM_IMAGES, g, &w.errNum));
 	int mem_b = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR,
                sizeof(int) * 256*NUM_IMAGES, b, &w.errNum));

	//set up the arguments for the kernel
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 0, sizeof(cl_mem), &output_image ),
                "Error setting kernel arguments10.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 1, sizeof(cl_mem), &my_sampler),
                "Error setting kernel arguments11.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 2, sizeof(cl_mem), &w.memObjects[mem_r]),
                "Error setting kernel arguments12.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 3, sizeof(cl_mem), &w.memObjects[mem_g]),
                "Error setting kernel arguments13.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 4, sizeof(cl_mem), &w.memObjects[mem_b]),
                "Error setting kernel arguments14.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 5, sizeof(cl_int), &single_image_width),
                "Error setting kernel arguments15.");
	 w.check(clSetKernelArg(w.kernels["build_color_histograms"], 6, sizeof(cl_int), &width),
                "Error setting kernel arguments16.");


        w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["build_color_histograms"], 1, NULL,
                hist_globalWorkSize, hist_localWorkSize, 0, NULL, &w.events[HISTOGRAM_TIME][event_number[HISTOGRAM_TIME]++]),
                "Error queuing kernel for execution.");
        
    //DO WE NEED TO READ BACK THESE VALUES?
    w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[mem_r], CL_TRUE, 0,
            256 * NUM_IMAGES * sizeof(int), r, 0, NULL, &w.events[HISTOGRAM_TIME][event_number[HISTOGRAM_TIME]++]),
            "Error reading to the result boarda.");
    w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[mem_g], CL_TRUE, 0,
            256 * NUM_IMAGES * sizeof(int), g, 0, NULL, &w.events[HISTOGRAM_TIME][event_number[HISTOGRAM_TIME]++]),
            "Error reading to the result boarbd.");
    w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[mem_b], CL_TRUE, 0,
            256 * NUM_IMAGES * sizeof(int), b, 0, NULL, &w.events[HISTOGRAM_TIME][event_number[HISTOGRAM_TIME]++]),
            "Error reading to the result boardc.");
//Uncomment for raw outout of the r histogram
/*
for (int i = 0; i < 4; i ++){
	for(int j = 0; j < 256; j++)
	{
		cout << r[j + i*256] << ", ";
	}
	cout << endl;
}    
*/
	/******************
  	 GET COLORS  Kernel
         WE CANNOT DO THIS AS WE DID IN THE HOST.  ALTERATIONS TO HOST+KERNEL CODE
	*******************/


	int num_colors = COLORS_PER_CHANNEL;	
	int* r_colors = (int*) malloc(sizeof(int) * NUM_IMAGES * num_colors);
	int* g_colors = (int*)malloc(sizeof(int) * NUM_IMAGES * num_colors);
	int* b_colors = (int*)malloc(sizeof(int) * NUM_IMAGES * num_colors);
	int pixels_per_color = (width*height) / (num_colors*NUM_IMAGES);
cout << "PPC: " << pixels_per_color << endl;
	for (int i=0;i<NUM_IMAGES*num_colors;i++){
		r_colors[i] = 0;
		g_colors[i] = 0;
		b_colors[i] = 0;
	}



        size_t color_globalWorkSize[1] =  { NUM_IMAGES } ;
        size_t color_localWorkSize[1] = { 1 } ;

    	int r_color_num = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE,
                sizeof(int) * NUM_IMAGES * num_colors, r_colors, &w.errNum));

    	int g_color_num = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE,
                sizeof(int) * NUM_IMAGES * num_colors, g_colors, &w.errNum));

    	int b_color_num = w.addMemObject(clCreateBuffer(w.context, CL_MEM_READ_WRITE,
                sizeof(int) * NUM_IMAGES * num_colors, b_colors, &w.errNum));


	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 0, sizeof(cl_mem), &w.memObjects[mem_r] ),
                "Error setting kernel arguments20.");

	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 1, sizeof(cl_mem), &w.memObjects[r_color_num] ),
                "Error setting kernel arguments21.");

	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 2, sizeof(cl_int),  &pixels_per_color),
                "Error setting kernel arguments22.");
	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 3, sizeof(cl_int), &num_colors ),
                "Error setting kernel arguments23.");


        w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["get_colors_from_histograms"], 1, NULL,
                color_globalWorkSize, color_localWorkSize, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]),
                "Error queuing kernel for execution.");

	w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[r_color_num], CL_TRUE, 0,
            num_colors * NUM_IMAGES * sizeof(int), r_colors, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]), 
	     "Error reading to the result board.");

	w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 0, sizeof(cl_mem), &w.memObjects[mem_g] ), 
		"Error setting kernel arguments20." );

	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 1, sizeof(cl_mem), &w.memObjects[g_color_num] ),
                "Error setting kernel arguments21.");

        w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["get_colors_from_histograms"], 1, NULL,
                color_globalWorkSize, color_localWorkSize, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]),
                "Error queuing kernel for execution.");
	w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[g_color_num], CL_TRUE, 0,
            num_colors * NUM_IMAGES * sizeof(int), g_colors, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]),
            "Error reading to the result board.");


	w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 0, sizeof(cl_mem), &w.memObjects[mem_b] ),
                "Error setting kernel arguments20.");

	 w.check(clSetKernelArg(w.kernels["get_colors_from_histograms"], 1, sizeof(cl_mem), &w.memObjects[b_color_num] ),
                "Error setting kernel arguments21.");

        w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["get_colors_from_histograms"], 1, NULL,
                color_globalWorkSize, color_localWorkSize, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]),
                "Error queuing kernel for execution.");
	w.check(clEnqueueReadBuffer(w.commandQueue, w.memObjects[b_color_num], CL_TRUE, 0,
            num_colors * NUM_IMAGES * sizeof(int), b_colors, 0, NULL, &w.events[FIND_COLORS_TIME][event_number[FIND_COLORS_TIME]++]),
            "Error reading to the result board.");

//Uncommnent for raw output of the color tables
/*
	for(int i =0; i<num_colors * NUM_IMAGES; i++){
		cout << g_colors[i] << ", ";
		if (((i+1) % num_colors) == 0) cout << endl; 
	}
*/
	/***********************
  	COLORIZE Kernel
        WE CANNOT DO THIS AS WE DID IN THE HOST.  ALTERATIONS TO HOST+KERNEL CODE
	************************/

	//section to write back the image
	cl_mem final_image = clCreateImage2D(w.context,
                                       CL_MEM_READ_WRITE,
                                       &clImageFormat,
                                       width,
                                       height,
                                       0,
                                       NULL,
                                       &errNum);




	
	char *bufferTwo = new char [width * height * 4];
	size_t origin[3] = { 0, 0, 0 };
	size_t region[3] = { width, height, 1};
	for (int x = 0;x<width*height*4;x++){
		bufferTwo[x] = 0;
	}


	my_sampler = clCreateSampler(w.context,
				     CL_FALSE,
				     CL_ADDRESS_CLAMP_TO_EDGE,
			             CL_FILTER_NEAREST,
				     &errNum
				    );

	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 0, sizeof(cl_mem), &output_image ),
                "Error setting kernel arguments20.");
	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 1, sizeof(cl_mem), &final_image ),
                "Error setting kernel arguments20.");

	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 2, sizeof(cl_sampler), &my_sampler ),
                "Error setting kernel arguments20.");

	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 3, sizeof(cl_mem),  &w.memObjects[r_color_num]),
                "Error setting kernel arguments20.");

	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 4, sizeof(cl_mem),  &w.memObjects[g_color_num]),
                "Error setting kernel arguments20.");

	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 5, sizeof(cl_mem),  &w.memObjects[b_color_num]),
                "Error setting kernel arguments20.");
	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 6, sizeof(cl_mem), &w.memObjects[my_edges] ),
                "Error setting kernel argument 6 in build_image_from_new_colors.");
	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 7, sizeof(cl_int),  &num_colors),
                "Error setting kernel arguments20.");


	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 8, sizeof(cl_int),  &single_image_width),
                "Error setting kernel arguments20.");
	w.check(clSetKernelArg(w.kernels["build_image_from_new_colors"], 9, sizeof(cl_int),  &width),
                "Error setting kernel arguments20.");


	size_t build_localWorkSize[1] = { WORK_GROUP_SIZE } ;
        size_t build_globalWorkSize[1] =  { RoundUp(build_localWorkSize[0],width*height) } ;
        

        w.check(clEnqueueNDRangeKernel(w.commandQueue, w.kernels["build_image_from_new_colors"], 1, NULL,
                build_globalWorkSize, build_localWorkSize, 0, NULL, &w.events[REPLACE_PIXEL_TIME][event_number[REPLACE_PIXEL_TIME]++]),
                "Error queuing kernel for execution.");

	errNum = clEnqueueReadImage(w.commandQueue, final_image, CL_TRUE,
                                origin, region, 0, 0, bufferTwo,
                                0, NULL, &w.events[REPLACE_PIXEL_TIME][event_number[REPLACE_PIXEL_TIME]++]);

	char newName[80] = "outputImages/big1.tif";
	
	if(!SaveImage(newName, bufferTwo, width,height)){
		std:cerr << "Error reading the result buffer or writing the results image." << endl;
	}
	else cout << "Successfully written to out" <<  newName << endl;
	delete[] bufferTwo;
	
	
	free(r_colors);
	free(g_colors);
	free(b_colors);	
	free(edge_array);
	//w.releaseMemObjects();
	if (clReleaseMemObject(output_image) != CL_SUCCESS) cout <<"release error 1" << endl;
	if (clReleaseMemObject(final_image) != CL_SUCCESS)  cout <<"release error 2" << endl;
	if (clReleaseMemObject(test_image) != CL_SUCCESS)  cout <<"release error 3" << endl;


    }
    catch (cl_error& e) {
        std::cerr << e.what() << " " << e.code << std::endl;
        w.cleanup();
        return 1;
    }

    cl_ulong totalTime[6] = {0,0,0,0,0,0};
    cl_ulong finalTime = 0;
    for (int h=0;h < 5;h++){
	
	    for (int i =0; i<event_number[h];i++)
	    {
		cl_ulong start, end;
		w.check(clGetEventProfilingInfo(w.events[h][i], CL_PROFILING_COMMAND_START,
		        sizeof(cl_ulong), &start, NULL),
		        "Error getting event profile information");
		w.check(clGetEventProfilingInfo(w.events[h][i], CL_PROFILING_COMMAND_END,
		        sizeof(cl_ulong), &end, NULL),
		        "Error getting event profile information");

		totalTime[h] += (end - start);
		finalTime += (end - start);
	    }
    }


    cout << "The total time for gaussians was: " << totalTime[GAUSSIAN_TIME] / 1000000000.0 << " seconds." << endl;
    cout << "The total time for sobels was: " << totalTime[SOBEL_TIME] / 1000000000.0 << " seconds." << endl;
    cout << "The total time for histograms was: " << totalTime[HISTOGRAM_TIME] / 1000000000.0 << " seconds." << endl;
    cout << "The total time for finding new colors was: " << totalTime[FIND_COLORS_TIME] / 1000000000.0 << " seconds." << endl;
    cout << "The total time for replacing pixels was: " << totalTime[REPLACE_PIXEL_TIME] / 1000000000.0 << " seconds." << endl;
    //the total time for everything
    cout << "The total time for everything was: " << finalTime / 1000000000.0 << " seconds." << endl;
    w.cleanup();

 
    
    return 0;
}
